Skip to content

cuTile smoke test: add sm_121 fragment for GB10, document a driver JIT-link gap - #2568

Open
zbrad wants to merge 1 commit into
NVIDIA:release/26.10from
zbrad:gb10-cutile-sm121-fix
Open

cuTile smoke test: add sm_121 fragment for GB10, document a driver JIT-link gap#2568
zbrad wants to merge 1 commit into
NVIDIA:release/26.10from
zbrad:gb10-cutile-sm121-fix

Conversation

@zbrad

@zbrad zbrad commented Sep 9, 2026

Copy link
Copy Markdown

Summary

GB10 (compute capability 12.1, DGX Spark) rejects cutile_smoke's only embedded cubin (sm_120) at cudaLibraryGetKernel with cudaErrorNoKernelImageForDevice, despite matching find_compatible_cubin_fragment()'s same-major/lower-minor forward-compatibility rule. Confirmed empirically on real GB10 hardware that this "forward compatible within a major family" assumption doesn't hold here.

  • Adds an exact-match cutile_arch_12_1 / sm_121 fragment.
  • sm_121a (the family-conditional target) isn't reachable yet — tileiras in the currently-published cuda-tile doesn't support any -a suffixed --gpu-name target (confirmed by testing sm_90a/sm_100a/sm_120a/sm_121a, all rejected identically); filed as NVIDIA/cutile-python#105.
  • Corrects find_compatible_cubin_fragment()'s doc comment, which stated cubins are forward compatible across minor revisions within a major family — not true for this device.

A second, independent finding

Fixing the above surfaced a separate problem: actually launching any cuTile-compiled kernel on this GB10 box fails with cudaErrorJitCompilerNotFound. Traced via strace to a real, confirmed cause: libnvidia-gpucomp.so, needed at kernel-launch time to complete a runtime relocation cuTile's cubins carry (.rela.nv.constant4 against ___NV_TILE_LAUNCH..., absent from a plain nvcc-compiled cubin through the identical cudaLibraryLoadData/cudaLibraryGetKernel/cudaLaunchKernelExC sequence — that path works fine), isn't shipped by the driver branch tested against (580.173.02).

Since is_expected_cutile_unavailable() already treats cudaErrorJitCompilerNotFound as an expected-unavailable condition for the load step, LaunchesCompatibleCubin now also tolerates it at launch time (skip, not fail) for consistency. A new CutileSmoke.RequiresJitLinkCapableDriver test deliberately does not tolerate it, so this gap stays visible rather than silently skipped — it's expected to start passing once a driver ships the missing component, which is the signal to fold its guarantee back into LaunchesCompatibleCubin and delete both.

Testing

  • CutileSmoke.ResolvesEveryEmbeddedArchitecture — updated for the new fragment, verified passing.
  • CutileSmoke.LaunchesCompatibleCubin — was failing (ASSERT_NE(launcher, nullptr)), now resolves the launcher correctly (confirms the arch fix); skips cleanly on this driver due to the second issue above.
  • CutileSmoke.RequiresJitLinkCapableDriver (new) — fails as expected/documented on this driver branch.
  • Rest of the GB10 gtest suite (34/35 binaries, ~8000+ tests) run clean against this change; the one other failure (CLUSTER_TEST, a KMeans++ random-subsample tolerance flake) is unrelated and reproducibly flaky independent of this change.

Assisted-by: Claude

@zbrad
zbrad requested a review from a team as a code owner September 9, 2026 02:08
Copilot AI lite review requested due to automatic review settings September 9, 2026 02:08
@copy-pr-bot

copy-pr-bot Bot commented Sep 9, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

It introduces a test explicitly documented as “expected to FAIL” (risking CI instability) and includes a helper that can leak CUDA allocations on ASSERT_* early-exit paths.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

This PR updates the cuTile smoke-test infrastructure to properly support GB10 (SM 12.1) by embedding an exact-match sm_121 cubin fragment, and it adjusts test behavior/documentation to reflect that “same-major/lower-minor” SASS fallback is not reliable for this family.

Changes:

  • Add cutile_arch_12_1 / sm_121 to the embedded cubin matrix and register it in the smoke-test fragment list/planner.
  • Update the find_compatible_cubin_fragment() doc comment to qualify the forward-compatibility assumption for SM 12.x (GB10).
  • Refactor the smoke launch test to centralize kernel execution/verification and add a new driver capability test (currently described as expected to fail on some drivers).
File summaries
File Description
cpp/tests/detail/jit_lto/cutile_smoke.cu Adds SM 12.1 fragment usage; refactors launch/verify path; adds a new driver JIT-link capability test.
cpp/src/detail/jit_lto/cutile_smoke/cutile_smoke_matrix.json Adds sm_121 cubin export entry for the smoke kernel.
cpp/include/cuvs/detail/jit_lto/cutile_module.hpp Updates compatibility-selection documentation to note SM 12.1 rejecting SM 12.0 SASS.
cpp/include/cuvs/detail/jit_lto/cutile_arch_tags.hpp Introduces cutile_arch_12_1 with documentation explaining why it’s required on GB10.
Review details
  • Files reviewed: 4/4 changed files
  • Comments generated: 4
  • Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment on lines +74 to +79
float* lhs = nullptr;
float* rhs = nullptr;
float* output = nullptr;
ASSERT_EQ(cudaMalloc(&lhs, sizeof(host_lhs)), cudaSuccess);
ASSERT_EQ(cudaMalloc(&rhs, sizeof(host_rhs)), cudaSuccess);
ASSERT_EQ(cudaMalloc(&output, sizeof(host_output)), cudaSuccess);
// upgrade ships the missing component, this test starts passing on its own,
// which is the signal to fold its guarantee back into LaunchesCompatibleCubin
// and delete this test and the skip branch above.
TEST(CutileSmoke, RequiresJitLinkCapableDriver)
Comment on lines +59 to +61
// Runs the add-two-tiles kernel and checks the result. Throws (via RTCX_CUDA_TRY,
// see rtcx/macros.hpp) on any CUDA failure -- callers decide whether to tolerate
// a specific expected failure or let it propagate as a test failure.
Comment on lines +175 to +179
// Deliberately NOT tolerant of cudaErrorJitCompilerNotFound, unlike
// LaunchesCompatibleCubin above -- this is expected to FAIL on any driver
// branch missing libnvidia-gpucomp.so (confirmed absent from the pinned
// 580.173.02 branch via strace; only present in an untested 595.45.04
// package cached locally, not installed). Excluded from tuned/full_test.sh's
@divyegala divyegala added bug Something isn't working non-breaking Introduces a non-breaking change labels Sep 9, 2026

@divyegala divyegala left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you for catching this bug! Could you please re-target to release/26.10 branch?

// upgrade ships the missing component, this test starts passing on its own,
// which is the signal to fold its guarantee back into LaunchesCompatibleCubin
// and delete this test and the skip branch above.
TEST(CutileSmoke, RequiresJitLinkCapableDriver)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This test needs to be opt-in, otherwise it will fail on machines that have an embedded cubin but their drivers are not JIT capable

@zbrad
zbrad changed the base branch from main to release/26.10 September 9, 2026 05:08
@zbrad
zbrad requested review from a team as code owners September 9, 2026 05:08
@zbrad
zbrad requested a review from bdice September 9, 2026 05:08
GB10 (compute capability 12.1) rejects the existing sm_120 embedded
cubin outright (cudaErrorNoKernelImageForDevice) despite matching
find_compatible_cubin_fragment()'s same-major/lower-minor forward-compat
rule -- confirmed empirically that rule doesn't hold for this device.
Add an exact-match cutile_arch_12_1/sm_121 fragment (sm_121a is not
reachable: NVIDIA/cutile-python#105) and correct find_compatible_cubin_fragment's
doc comment, which overclaimed universal forward compatibility.

That surfaced a second, independent problem: actually launching any
cuTile-compiled kernel needs libnvidia-gpucomp.so at kernel-launch time
to complete a runtime relocation (confirmed via strace), and that
library genuinely isn't shipped by the currently-pinned 580.173.02
driver branch (only found, unused, in a cached 595.45.04 package).
LaunchesCompatibleCubin now tolerates cudaErrorJitCompilerNotFound as an
expected-unavailable skip during launch, mirroring what
is_expected_cutile_unavailable() already does for the load step. A new
CutileSmoke.DISABLED_RequiresJitLinkCapableDriver canary does NOT
tolerate it and is expected to fail until a future driver ships the
missing component -- DISABLED_ (GoogleTest's own opt-in convention) so
it doesn't run by default, only via --gtest_also_run_disabled_tests; its
own passing becomes the signal to fold the guarantee back into
LaunchesCompatibleCubin and delete both the canary and the skip branch.

Device buffers in the shared launch/verify helper are now held by
unique_ptr<float, CudaDeviceDeleter> so they're freed on every exit path
(including an early ASSERT_* return, which doesn't unwind via a C++
exception) rather than only on the success path.
@zbrad
zbrad force-pushed the gb10-cutile-sm121-fix branch from a2a79d4 to 30937fb Compare September 9, 2026 05:11
@zbrad

zbrad commented Sep 9, 2026

Copy link
Copy Markdown
Author

Thanks both — addressed everything, force-pushed:

  • Re-targeted to release/26.10 and reset the branch onto it directly (the diff is a clean 4-file/155-line change again, not the 51-file diff a bare base-branch swap produced against main's divergence).
  • Made the driver-gap test opt-in: renamed to DISABLED_RequiresJitLinkCapableDriver (GoogleTest's own convention — excluded from any default ctest/gtest run, only runs via --gtest_also_run_disabled_tests). It no longer relies on anything outside this repo to stay non-gating.
  • Fixed the leak: run_smoke_add_and_verify's three device buffers are now held by unique_ptr<float, CudaDeviceDeleter>, so they're freed on every exit path — including an early ASSERT_* return, which doesn't unwind via a C++ exception and was leaking them before.
  • Fixed the misleading doc comment: it now accurately describes that host-observable CUDA calls use ASSERT_* (early return, not a throw) while the kernel dispatch call itself throws via the vendored rtcx dependency — and clarifies rtcx/macros.hpp lives in that CPM-fetched dependency, not this repo's own tree.
  • Removed the tuned/full_test.sh reference — that was fork-only tooling not present here; the comment now just describes the DISABLED_ mechanism itself.

Verified with a real rebuild + compute-sanitizer --tool memcheck on GB10 hardware: default run is clean (DISABLED_RequiresJitLinkCapableDriver correctly excluded, no failures), explicit --gtest_also_run_disabled_tests --gtest_filter=*RequiresJitLinkCapableDriver run still fails exactly as documented, and the RAII change introduces no new sanitizer findings (the only two flagged events are the already-expected cudaErrorJitCompilerNotFound calls the test itself is built around).

zbrad added a commit to zbrad/cuvs that referenced this pull request Sep 9, 2026
Sync the same fixes from the upstream PR after real review feedback
(divyegala + Copilot bot):
- Fix a real leak: run_smoke_add_and_verify's device buffers are now
  held by unique_ptr<float, CudaDeviceDeleter>, freed on every exit
  path (an early ASSERT_* return doesn't unwind via a C++ exception and
  was leaking them before).
- Fix a misleading doc comment claiming the helper "throws... on any
  CUDA failure" when most of its CUDA calls use ASSERT_* (early return,
  not a throw) instead.
- Rename the driver-gap canary to DISABLED_RequiresJitLinkCapableDriver
  (GoogleTest's own opt-in convention), which supersedes this repo's
  custom --gtest_filter exclusion in tuned/full_test.sh -- simplified
  that script accordingly, and it no longer relies on fork-only tooling
  to stay non-gating.

Verified: default run clean (DISABLED_ test excluded, no failures),
--gtest_also_run_disabled_tests run still fails exactly as documented,
compute-sanitizer clean (only the two already-expected
cudaErrorJitCompilerNotFound events, no new findings from the RAII
change).
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

bug Something isn't working non-breaking Introduces a non-breaking change

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants